__device__ unsigned int
RLE_Encode_MN5(const unsigned char * /*[in]*/data, unsigned int /*[in]*/bytes, unsigned char * /*[out]*/encoded, unsigned int & /*[out]*/enc_bytes)
{
    // RLE (MN5 variation of length 4) encoding
    //  - example: "AAAAAAABBBBCCCD" -->  "AAAA\3BBBB\0CCCD"
    enc_bytes = 0;
    unsigned char c, n;
    int reps = 0;
    if(bytes > 0)
    {
        c = encoded[enc_bytes++] = data[0];
        for(unsigned int i = 1; i < bytes; i++)
        {
            if(data[i] == c)
            {
                ++reps;
                if(reps <= 4)
                    encoded[enc_bytes++] = c;
                else
                {
                    if(reps > 255)	// == 256
                    {
                        n = reps - 4 - 1;
                        encoded[enc_bytes++] = n;
                        encoded[enc_bytes++] = c;
                        reps = 1;
                    }
                }
            }
            else
            {
                if(reps >= 4)
                {
                    n = reps - 4;
                    encoded[enc_bytes++] = n;
                }
                reps = 1;
                c = data[i];
                encoded[enc_bytes++] = c;
            }
                
        }
    }
    if(reps >= 4) encoded[enc_bytes++] = unsigned char(reps - 4);
    //
    return enc_bytes;
}

__device__ unsigned int
RLE_Decode_MN5(const unsigned char * /*[in]*/data, unsigned int /*[in]*/bytes, unsigned char * /*[out]*/decoded, unsigned int & /*[out]*/dec_bytes)
{
    // RLE (MN5 variation of length 4) decoding
    //  - example: "AAAA\3BBBB\0CCCD" --> "AAAAAAABBBBCCCD"
    dec_bytes = 0;
    unsigned char c;
    int reps = 1;
    if(bytes > 0)
    {
        c = decoded[dec_bytes++] = data[0];
        for(unsigned int i = 1; i < bytes; i++)
        {
            // next character is a number of repetitions
            if(reps == 4)
            {
                for(int r = data[i]; r > 0; r--)
                    decoded[dec_bytes++] = c;
                reps = 1;
                if((++i) < bytes)
                    decoded[dec_bytes++] = c = data[i];
                continue;
            }
            // a character
            if(data[i] == c)
                ++reps;
            else
            {
                reps = 1;
                c = data[i];
            }
            decoded[dec_bytes++] = c;
        }
    }
    //
    return dec_bytes;
}